/**
 * 
 */
package sw.opencl;

import static org.jocl.CL.CL_CONTEXT_PLATFORM;
import static org.jocl.CL.CL_DEVICE_NAME;
import static org.jocl.CL.CL_MEM_READ_ONLY;
import static org.jocl.CL.CL_PLATFORM_NAME;
import static org.jocl.CL.CL_TRUE;
import static org.jocl.CL.clBuildProgram;
import static org.jocl.CL.clCreateBuffer;
import static org.jocl.CL.clCreateCommandQueue;
import static org.jocl.CL.clCreateContext;
import static org.jocl.CL.clCreateKernel;
import static org.jocl.CL.clCreateProgramWithSource;
import static org.jocl.CL.clEnqueueNDRangeKernel;
import static org.jocl.CL.clEnqueueReadBuffer;
import static org.jocl.CL.clGetDeviceIDs;
import static org.jocl.CL.clGetPlatformIDs;
import static org.jocl.CL.clReleaseCommandQueue;
import static org.jocl.CL.clReleaseContext;
import static org.jocl.CL.clReleaseEvent;
import static org.jocl.CL.clReleaseKernel;
import static org.jocl.CL.clReleaseMemObject;
import static org.jocl.CL.clReleaseProgram;
import static org.jocl.CL.clSetKernelArg;
import static org.jocl.CL.clWaitForEvents;
import static org.jocl.CL.clGetDeviceInfo;

import java.io.File;
import java.io.IOException;
import java.nio.ByteBuffer;
import java.nio.ByteOrder;

import org.jocl.CL;
import org.jocl.Pointer;
import org.jocl.Sizeof;
import org.jocl.cl_command_queue;
import org.jocl.cl_context;
import org.jocl.cl_context_properties;
import org.jocl.cl_device_id;
import org.jocl.cl_event;
import org.jocl.cl_kernel;
import org.jocl.cl_mem;
import org.jocl.cl_platform_id;
import org.jocl.cl_program;

import sw.utils.Direction;
import sw.utils.JOCLUtils;

/**
 * @author Chung
 *
 */
public class SW_OpenCL_V3_MultipleDevices {
	
	private final int GAP = -1;
	private int match = 2;
	private int mismatch = -1;
	private static String KERNEL_FILE_NAME = "/kernel_v3.cl";
	
	private void sw(String s1, String s2, final int sub) {
		int col = s1.length(); // column
		int row = s2.length(); // row
		int size = (row+1)*(col+1);
		int matrix[] = new int[size];
		
		String path = new File("").getAbsolutePath();
		String kernelSrc = "";
		try {
			kernelSrc = JOCLUtils.loadKernel(path + KERNEL_FILE_NAME);
//			System.out.println(kernelSrc);
		} catch (IOException e) {
			// TODO Auto-generated catch block			
			e.printStackTrace();
			return;
		}

		int newRow = row / sub;
		int newCol = col / sub;

//		The platform and device type that will be used
		final int platformIndex = 0;
		final long deviceType = CL.CL_DEVICE_TYPE_ALL;   
		final int deviceIndex = 1; // CPU

       // Enable exceptions and subsequently omit error checks in this sample
		CL.setExceptionsEnabled(true);

       // Obtain the number of platforms
		int numPlatformsArray[] = new int[1];
		clGetPlatformIDs(0, null, numPlatformsArray);
		int numPlatforms = numPlatformsArray[0];

		//Obtain a platform ID
		cl_platform_id platforms[] = new cl_platform_id[numPlatforms];
		clGetPlatformIDs(platforms.length, platforms, null);
		cl_platform_id platform = platforms[platformIndex];
//		String platformName = JOCLUtils.getString(platform, CL_PLATFORM_NAME);
//		System.out.println("Using platform "+platformIndex+" of "+ numPlatforms+": "+platformName);

        // Initialize the context properties
		cl_context_properties contextProperties = new cl_context_properties();
		contextProperties.addProperty(CL_CONTEXT_PLATFORM, platform);
       
       // Obtain the number of devices for the platform
		int numDevicesArray[] = new int[1];
		clGetDeviceIDs(platform, deviceType, 0, null, numDevicesArray);
		int numDevices = numDevicesArray[0];
       
       // Obtain a device IDs 
		cl_device_id devices[] = new cl_device_id[numDevices];
		clGetDeviceIDs(platform, deviceType, numDevices, devices, null);
		
		//int maxWorkGroupSize = getMaxWorkGroupSize(devices[deviceIndex]);
		//System.out.println("maxWorkGroupSize " + maxWorkGroupSize);

//		int subMatrix = (int) Math.sqrt(maxWorkGroupSize);
//		System.out.println("submatrix size " + subMatrix);
//		int newRow = row / subMatrix;
//		int newCol = col / subMatrix;
		
	//	String deviceName = JOCLUtils.getString(devices[deviceIndex], CL_DEVICE_NAME);
	//	System.out.println("Device :" + deviceName);

       // Create a context for the devices
		cl_context context = clCreateContext(contextProperties, 1, new cl_device_id[]{devices[deviceIndex]}, null, null, null);

       // Create and build the program and the kernel
       cl_program program = clCreateProgramWithSource(context, 1, new String[]{ kernelSrc }, null, null);
       clBuildProgram(program, 0, null, null, null, null);
       cl_kernel kernel = clCreateKernel(program, "sw", null);

       // Allocate the memory objects for the input- and output data
       cl_mem matrixMem = clCreateBuffer(context, CL.CL_MEM_READ_WRITE | CL.CL_MEM_COPY_HOST_PTR,  
									  Sizeof.cl_int * matrix.length, Pointer.to(matrix), null);		
		
       cl_mem dimMem = clCreateBuffer(context, CL.CL_MEM_READ_ONLY | CL.CL_MEM_COPY_HOST_PTR,  
				  Sizeof.cl_int, Pointer.to(new int[]{col+1}), null);	
		
       cl_mem s1Mem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL.CL_MEM_COPY_HOST_PTR,  
				  Sizeof.cl_char * (s1.length()), Pointer.to(s1.getBytes()), null);
		
       cl_mem s2Mem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL.CL_MEM_COPY_HOST_PTR , 
				  Sizeof.cl_char * (s2.length()), Pointer.to(s2.getBytes()), null);
		
       // Create command-queue for device
       long properties = 0;
       properties |= CL.CL_QUEUE_PROFILING_ENABLE ;
       cl_command_queue commandQueue = clCreateCommandQueue(context, devices[deviceIndex], properties, null);
       
       // Execute the kernel on each command queue, and 
       // create events for each kernel launch
       clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(matrixMem));
       clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(dimMem));		
       clSetKernelArg(kernel, 4, Sizeof.cl_mem, Pointer.to(s1Mem));
       clSetKernelArg(kernel, 5, Sizeof.cl_mem, Pointer.to(s2Mem));
       
       long before = System.nanoTime();
       for (int line = 1; line <= (newRow + newCol - 1); line++) {
			/*
			 * Get column index of the first element in this line of output. The
			 * index is 0 for first ROW lines and line - ROW for remaining lines
			 */
			int start_col = Math.max(0, line - newRow);

			// number of diagonals of matrix
			int subLength = JOCLUtils.min(line, (newCol - start_col), newRow);
			int subIndexes[] = new int[2 * subLength];

			/* store coordinates of cells on the same diagonal */
			for (int j = 0; j < subLength; j++) {
				subIndexes[2 * j] = (Math.min(newRow, line) - j - 1) * sub;
				subIndexes[2 * j + 1] = (start_col + j) * sub;
			}

			cl_mem subIndexMem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL.CL_MEM_COPY_HOST_PTR, 
												Sizeof.cl_int * (subIndexes.length), Pointer.to(subIndexes), null);

			cl_mem subLengthMem = clCreateBuffer(context, CL_MEM_READ_ONLY | CL.CL_MEM_COPY_HOST_PTR, 
												Sizeof.cl_int, Pointer.to(new int[] { sub }), null);

			clSetKernelArg(kernel, 2, Sizeof.cl_mem, Pointer.to(subIndexMem));
			clSetKernelArg(kernel, 3, Sizeof.cl_mem, Pointer.to(subLengthMem));

			cl_event event = new cl_event();
			clEnqueueNDRangeKernel(commandQueue, kernel, 1, null, new long[] { (long) (subLength) }, null, 0, null, event);
			// waiting for kernel completed
			clWaitForEvents(1, new cl_event[] { event });
			clReleaseEvent(event);
			clReleaseMemObject(subIndexMem);
			clReleaseMemObject(subLengthMem);
       }
		
		 // Read the output data of the first device
        clEnqueueReadBuffer(commandQueue, matrixMem, CL_TRUE, 0, size * Sizeof.cl_int, 
        				Pointer.to(matrix), 0, null, null);
        long after = System.nanoTime();
		float totalDurationMs = (after-before)/1e6f;
		System.out.println("Total duration: "+totalDurationMs+"ms");
        
        // Release kernel, program, and memory objects
        clReleaseKernel(kernel);
        clReleaseProgram(program);
        clReleaseMemObject(matrixMem);
        clReleaseMemObject(dimMem);
        clReleaseMemObject(s1Mem);
        clReleaseMemObject(s2Mem);
        clReleaseCommandQueue(commandQueue);
        clReleaseContext(context);
	}
	
	private void printMatrix(int[] matrix, int row, int col) {
		for (int i = 0; i < ((row+1)*(col+1)); i++) {			
			System.out.print(matrix[i] + " ");
			if ( (i+1) % (col+1) == 0) {
				System.out.println("");
			}
		}
	}
	
	int getDiagonalLength(int slice, int z, int sub) {
		  int i = z;
		  int j = slice - z;

		  if (z <= 0)
		    return  j + 1;

		  return  j - i + 1;
		}
	
	
	
	private static int getMaxWorkGroupSize(cl_device_id device) {
	        ByteBuffer buffer = ByteBuffer.allocate(Sizeof.size_t).order(ByteOrder.nativeOrder());
	        clGetDeviceInfo(device, CL.CL_DEVICE_MAX_WORK_GROUP_SIZE, Sizeof.size_t, Pointer.to(buffer), null);
	        if (Sizeof.size_t == 4) {
	            return buffer.getInt();
	        } else {
	            return (int)buffer.getLong();
	        }
	    }
	
	/**
	 *  args
	 */
	public static void main(String[] args) {		
		
		SW_OpenCL_V3_MultipleDevices swOpenCL = new SW_OpenCL_V3_MultipleDevices();
		String s1 = JOCLUtils.generateSequenceStringWithLength(1000);
		String s2 = JOCLUtils.generateSequenceStringWithLength(1000);
        int subMatrix = JOCLUtils.getGreatestDivisor(s2.length(), s1.length());
//        newRow = row / sub;
//        newCol = col / sub;
//        System.out.println(JOCLUtils.getGreatestDivisor(1000, 600));
//        long before = System.nanoTime();
		swOpenCL.sw(s1, s2, 50);
		
//		long after = System.nanoTime();
//		float totalDurationMs = (after-before)/1e6f;
//		System.out.println("Total duration: "+totalDurationMs+"ms");
	}
	
}
